#include "KDop.h"
#include "..//Scene//Scene.h"
#include <cuda_runtime.h>

Tuple3f SplitPlane[3] = 
{
	Tuple3f(1.0f, 0.0f, 0.0f),
	Tuple3f(0.0f, 1.0f, 0.0f),
	Tuple3f(0.0f, 0.0f, 1.0f),
};

#define MaxTrianglesPerLeaf 5
#define Max1DTextureSize 65536
//Vertex index buffer textures
#define MaxVertexIndexBufferTexturesCount 10
texture<int, 1, cudaReadModeElementType>* pVertexIndexBufferTextures[MaxVertexIndexBufferTexturesCount];

#define MaxTriangleIndexBufferTexturesCount 10
texture<int, 1, cudaReadModeElementType>* pTriangleIndexBufferTextures[MaxTriangleIndexBufferTexturesCount];

texture<int, 1, cudaReadModeElementType> VertexIndexBufferTextures_0;
texture<int, 1, cudaReadModeElementType> VertexIndexBufferTextures_1;
texture<int, 1, cudaReadModeElementType> VertexIndexBufferTextures_2;
texture<int, 1, cudaReadModeElementType> VertexIndexBufferTextures_3;
texture<int, 1, cudaReadModeElementType> VertexIndexBufferTextures_4;
texture<int, 1, cudaReadModeElementType> VertexIndexBufferTextures_5;
texture<int, 1, cudaReadModeElementType> VertexIndexBufferTextures_6;
texture<int, 1, cudaReadModeElementType> VertexIndexBufferTextures_7;
texture<int, 1, cudaReadModeElementType> VertexIndexBufferTextures_8;
texture<int, 1, cudaReadModeElementType> VertexIndexBufferTextures_9;

texture<int, 1, cudaReadModeElementType> TriangleIndexBufferTextures_0;
texture<int, 1, cudaReadModeElementType> TriangleIndexBufferTextures_1;
texture<int, 1, cudaReadModeElementType> TriangleIndexBufferTextures_2;
texture<int, 1, cudaReadModeElementType> TriangleIndexBufferTextures_3;
texture<int, 1, cudaReadModeElementType> TriangleIndexBufferTextures_4;
texture<int, 1, cudaReadModeElementType> TriangleIndexBufferTextures_5;
texture<int, 1, cudaReadModeElementType> TriangleIndexBufferTextures_6;
texture<int, 1, cudaReadModeElementType> TriangleIndexBufferTextures_7;
texture<int, 1, cudaReadModeElementType> TriangleIndexBufferTextures_8;
texture<int, 1, cudaReadModeElementType> TriangleIndexBufferTextures_9;

_DEVICE_ONLY_ int GetVertexIndexValue(int Index)
{
	int iTextureIndex = Index / Max1DTextureSize;
	int iSampleIndex = Index - iTextureIndex * Max1DTextureSize;
	switch(iTextureIndex)
	{
	case 0:
		return tex1Dfetch(VertexIndexBufferTextures_0, iSampleIndex);
	case 1:
		return tex1Dfetch(VertexIndexBufferTextures_1, iSampleIndex);
	case 2:
		return tex1Dfetch(VertexIndexBufferTextures_2, iSampleIndex);
	case 3:
		return tex1Dfetch(VertexIndexBufferTextures_3, iSampleIndex);
	case 4:
		return tex1Dfetch(VertexIndexBufferTextures_4, iSampleIndex);
	case 5:
		return tex1Dfetch(VertexIndexBufferTextures_5, iSampleIndex);
	case 6:
		return tex1Dfetch(VertexIndexBufferTextures_6, iSampleIndex);
	case 7:
		return tex1Dfetch(VertexIndexBufferTextures_7, iSampleIndex);
	case 8:
		return tex1Dfetch(VertexIndexBufferTextures_8, iSampleIndex);
	case 9:
		return tex1Dfetch(VertexIndexBufferTextures_9, iSampleIndex);
	}
	return 0;
	//return tex1Dfetch(VertexIndexBufferTextures[iTextureIndex], iSampleIndex);
}

_DEVICE_ONLY_ int GetTriangleIndexValue(int Index)
{
	int iTextureIndex = Index / Max1DTextureSize;
	int iSampleIndex = Index - iTextureIndex * Max1DTextureSize;
	switch(iTextureIndex)
	{
	case 0:
		return tex1Dfetch(TriangleIndexBufferTextures_0, iSampleIndex);
	case 1:
		return tex1Dfetch(TriangleIndexBufferTextures_1, iSampleIndex);
	case 2:
		return tex1Dfetch(TriangleIndexBufferTextures_2, iSampleIndex);
	case 3:
		return tex1Dfetch(TriangleIndexBufferTextures_3, iSampleIndex);
	case 4:
		return tex1Dfetch(TriangleIndexBufferTextures_4, iSampleIndex);
	case 5:
		return tex1Dfetch(TriangleIndexBufferTextures_5, iSampleIndex);
	case 6:
		return tex1Dfetch(TriangleIndexBufferTextures_6, iSampleIndex);
	case 7:
		return tex1Dfetch(TriangleIndexBufferTextures_7, iSampleIndex);
	case 8:
		return tex1Dfetch(TriangleIndexBufferTextures_8, iSampleIndex);
	case 9:
		return tex1Dfetch(TriangleIndexBufferTextures_9, iSampleIndex);
	}
	return 0;
	//return tex1Dfetch(TriangleIndexBufferTextures[iTextureIndex], iSampleIndex);
}

//Triangle index buffer textures
//check 
struct ConstInfoIndex
{
	int iRenderConfigOffset;
	int iMeshInfoOffset;
	int iLightInfoOffset;
	int iSampleDirOffset;
	_HOST_DEVICE_ ConstInfoIndex()
	{
		iRenderConfigOffset = 0;
		iMeshInfoOffset = 0;
		iLightInfoOffset = 0;
		iSampleDirOffset = 0;
	}
};

//whole const buffer will be split to many slice for struct
#define MaxConstantBufferSize 60*1024
__constant__ char dConstantBuffer[MaxConstantBufferSize];
int hConstantBuffer[MaxConstantBufferSize/sizeof(int)];

_DEVICE_ONLY_ float GetHitTriangelUVColor(int iTrangleIndex
										, Tuple2f HittedUV
										, int* pConstantData
										, int Offset)
{
	int StaticMeshCount = pConstantData[Offset];
	int iCurOffset = Offset+sizeof(int)/sizeof(int);
	float HittedValue = 0.0f;
	for(int i=0; i<StaticMeshCount; ++i)
	{
		Tuple2i StaticMeshVertexInfo = *(Tuple2i*)(pConstantData+iCurOffset);
		iCurOffset += sizeof(Tuple2i)/sizeof(int);
		Tuple2i StaticMeshTextureSize = *(Tuple2i*)(pConstantData+iCurOffset);
		iCurOffset += sizeof(Tuple2i)/sizeof(int);
		unsigned char* pRaster = (unsigned char*)(pConstantData+iCurOffset);
		iCurOffset += sizeof(unsigned char*)/sizeof(int);
		//we find the hitted information
		if(iTrangleIndex>=StaticMeshVertexInfo.x && iTrangleIndex<StaticMeshVertexInfo.y)
		{
			int SamplePosX = HittedUV.y * StaticMeshTextureSize.y;
			int SamplePosY = HittedUV.x * StaticMeshTextureSize.x;
			SamplePosX = ClampT(SamplePosX, 0, StaticMeshTextureSize.x-1);
			SamplePosY = ClampT(SamplePosY, 0, StaticMeshTextureSize.y-1);
			unsigned char BufferValue = pRaster[SamplePosY*StaticMeshTextureSize.x+SamplePosX];
			HittedValue = BufferValue/float(255.0f);
			break;
		}
	}
	return HittedValue;
}

_DEVICE_ONLY_  void		ProcessIndirectionalLighting(Tuple3f& WorldPos
												  , Tuple3f& WorldNormal
												  , CVertex* pVertexBuffer
												 // , int* pVertexIndexBuffer
												 //, int* pTriangleIndexBuffer
												  , CKDopNode* pKDTreeRoot
												  , int* pConstantData
												  , unsigned char* pColor)
{
	ConstInfoIndex pInfoIndex = *(ConstInfoIndex*)pConstantData;
	CRenderConfig RenderConfig = *(CRenderConfig*)(pConstantData+pInfoIndex.iRenderConfigOffset);
	float fIrrandiance = 0.0f;

	CGPUStack<KDNodePair> LocalStack;
	Tuple3f UVT;
	Tuple3f CurSampleDir;
	float cosValue;
	/*for(int k=0; k<400; ++k)
	{*/
		for(int j=0; j<RenderConfig.SampleDirCount; j++)
		{
			CurSampleDir = *(Tuple3f*)(pConstantData+pInfoIndex.iSampleDirOffset+j*sizeof(Tuple3f)/sizeof(int));
			//check if upper
			cosValue = CurSampleDir.getDotProduct(WorldNormal);
			if(cosValue>0)
			{
				int iTrangleIndex = SceneLineCheck(WorldPos
													, CurSampleDir
													, UVT
													, pVertexBuffer
													//, pVertexIndexBuffer
													//, pTriangleIndexBuffer
													, pKDTreeRoot
													, LocalStack);
				if(iTrangleIndex>=0)
				{
					//int3 iCurTexelTriangleIdx = *((int3*)(&pVertexIndexBuffer[pTriangleIndexBuffer[iTrangleIndex]*3]));

					int3 iCurTexelTriangleIdx;
					iCurTexelTriangleIdx.x = GetVertexIndexValue(GetTriangleIndexValue(iTrangleIndex)*3);
					iCurTexelTriangleIdx.y = GetVertexIndexValue(GetTriangleIndexValue(iTrangleIndex)*3+1);
					iCurTexelTriangleIdx.z = GetVertexIndexValue(GetTriangleIndexValue(iTrangleIndex)*3+2);

					Tuple2f HittedUV = pVertexBuffer[iCurTexelTriangleIdx.x].UV*UVT.x
									+ pVertexBuffer[iCurTexelTriangleIdx.y].UV*UVT.y
									+ pVertexBuffer[iCurTexelTriangleIdx.z].UV*(1.0f-UVT.x-UVT.y);
					HittedUV.clamp(0.0f, 1.0f);
					fIrrandiance += GetHitTriangelUVColor(iTrangleIndex, HittedUV, pConstantData, pInfoIndex.iMeshInfoOffset)*cosValue;
				}
			}
		}
	//}
	fIrrandiance/=float(RenderConfig.SampleDirCount);
	int iColor = *pColor + fIrrandiance*256.0f;
	iColor = ClampT(iColor, 0, 255);
	*pColor = iColor;
}



_DEVICE_ONLY_  float ProcessDirectionalLightShadow(Tuple3f& WorldPos
												, Tuple3f& WorldNormal
												, CDirectionalLight* pLight
												, CVertex* pVertexBuffer
												//, int* pVertexIndexBuffer
												//, int* pTriangleIndexBuffer
												, CKDopNode* pKDTreeRoot)
{
	Tuple3f CheckDir(-pLight->Dir.x, -pLight->Dir.y, -pLight->Dir.z);
	float Brightness = ClampT(WorldNormal.getDotProduct(CheckDir), 0.0f, 1.0f);
	Tuple3f UVT(0, 0, -1.0f);
	CGPUStack<KDNodePair> LocalStack;
	if(SceneLineCheck(WorldPos
					, CheckDir
					, UVT
					, pVertexBuffer
					//, pVertexIndexBuffer
					//, pTriangleIndexBuffer
					, pKDTreeRoot
					, LocalStack)>=0 && UVT.z>0.0f )
	{
		Brightness = 0;
	}
	return Brightness;
}

_DEVICE_ONLY_  float ProcessPointLightShadow(Tuple3f& WorldPos
												   , Tuple3f& WorldNormal
												   , CPointLight* pLight
												   , CVertex* pVertexBuffer
												  // , int* pVertexIndexBuffer
												   //, int* pTriangleIndexBuffer
												   , CKDopNode* pKDTreeRoot)
{
	Tuple3f CheckDir = pLight->Position - WorldPos;
	float DistanceToLight = CheckDir.getLength();
	//CheckDir = CheckDir.normalize();
	CheckDir.x /= DistanceToLight;
	CheckDir.y /= DistanceToLight;
	CheckDir.z /= DistanceToLight;
	float Brightness = WorldNormal.getDotProduct(CheckDir);
	Tuple3f UVT(0, 0, FLT_MAX);
	CGPUStack<KDNodePair> LocalStack;
	if(SceneLineCheck(WorldPos
		, CheckDir
		, UVT
		, pVertexBuffer
		//, pVertexIndexBuffer
		//, pTriangleIndexBuffer
		, pKDTreeRoot
		, LocalStack)>=0)
	{
		//if(DistanceToLight>UVT.z)
		{
			Brightness = 0.0f;
		}
		/*Tuple3f HitDest = WorldPos + CheckDir*UVT.z;
		HitDest -= pLight->Position;
		if(HitDest.getLength()>pLight->SphereSize)
		{
			Brightness = 0;
		}*/
		//Brightness = 0;
	}
	//Brightness = ClampT(-UVT.z, 0.0f, 1.0f);
	return Brightness;
}

_DEVICE_ONLY_ float ProcessSpotLightShadow(Tuple3f& WorldPos
										   , Tuple3f& WorldNormal
										   , CSpotLight* pLight
										   , CVertex* pVertexBuffer
										  // , int* pVertexIndexBuffer
										  // , int* pTriangleIndexBuffer
										   , CKDopNode* pKDTreeRoot)
{
	return 0;
}

_DEVICE_ONLY_ float ProcessLightShadow(Tuple3f& WorldPos
									   , Tuple3f& WorldNormal
									   , CVertex* pVertexBuffer
									   //, int* pVertexIndexBuffer
									   //, int* pTriangleIndexBuffer
									   , CKDopNode* pKDTreeRoot
									   , int* pConstantData)
{
	ConstInfoIndex* pInfoIndex = (ConstInfoIndex*)pConstantData;
	Tuple3i LightCount = *(Tuple3i*)(pConstantData+pInfoIndex->iLightInfoOffset);
	int iCurOffset = sizeof(Tuple3i)/sizeof(int);
	float fRes = 0.0f;
	//process directional light
	for (int i=0; i<LightCount.x; ++i)
	{
		CDirectionalLight* pLight = (CDirectionalLight*)(pConstantData+pInfoIndex->iLightInfoOffset+iCurOffset);

		fRes += ProcessDirectionalLightShadow(WorldPos
											,WorldNormal
											,pLight
											,pVertexBuffer
											//,pVertexIndexBuffer
											//,pTriangleIndexBuffer
											,pKDTreeRoot);

		iCurOffset+=sizeof(CDirectionalLight)/sizeof(int);
	}
	//process point light
	for (int i=0; i<LightCount.y; ++i)
	{
		CPointLight* pLight = (CPointLight*)(pConstantData+pInfoIndex->iLightInfoOffset+iCurOffset);
		fRes += ProcessPointLightShadow(WorldPos
										,WorldNormal
										,pLight
										,pVertexBuffer
										//,pVertexIndexBuffer
										//,pTriangleIndexBuffer
										,pKDTreeRoot);
		iCurOffset+=sizeof(CPointLight)/sizeof(int);
	}
	//process spot light
	for (int i=0; i<LightCount.z; ++i)
	{
		CSpotLight* pLight = (CSpotLight*)(pConstantData+pInfoIndex->iLightInfoOffset+iCurOffset);
		fRes += ProcessSpotLightShadow(WorldPos
										,WorldNormal
										,pLight
										,pVertexBuffer
										//,pVertexIndexBuffer
										//,pTriangleIndexBuffer
										,pKDTreeRoot);	
		iCurOffset+=sizeof(CSpotLight)/sizeof(int);
	}
	return fRes;
}

_DEVICE_ONLY_  int SceneLineCheck(Tuple3f& Start
								, Tuple3f& Dir
								, Tuple3f& UVT
								, CVertex* pVertexBuffer
								//, int* pVertexIndexBuffer
								//, int* pTriangleIndexBuffer
								, CKDopNode* pKDTreeRoot
								, CGPUStack<KDNodePair>& LocalStack)
{
	//CGPUStack<KDNodePair> LocalStack;
	int		iHitTriangleIndex = -1;
	Tuple3f NearestHitInfo(0, 0, FLT_MAX);

	LocalStack.Clear();
	KDNodePair RootNode;
	LocalStack.PushElement(RootNode);

	const float fTollerance = 0.0001f;
	Tuple3f InvDir(1.0f/(Dir.x+fTollerance), 1.0f/(Dir.y+fTollerance) , 1.0f/(Dir.z+fTollerance));
	KDNodePair CurProcessNode;
	CKDopNode pCurNode;
	Tuple3f CurNearestHitInfo(0, 0, FLT_MAX);
	Tuple3i	iVertexIndexBuffer;
	float CurHitTime;
	float CurHitTimeLeft, CurHitTimeRight;
	bool bHit = true;

	while(!LocalStack.IsEmpty())
	{
		CurProcessNode = LocalStack.PopTop();
		pCurNode = pKDTreeRoot[CurProcessNode.NodeIndex];
		//check all triangles in this node
		if(pCurNode.IsLeaf)
		{
			for(int i=pCurNode.t.StartTriangleIndex; i<pCurNode.t.StartTriangleIndex+pCurNode.t.TriangleCount; ++i)
			{	
				//int		iCurVertexStartIndex = pTriangleIndexBuffer[i]*3;

				//iVertexIndexBuffer.x = pVertexIndexBuffer[iCurVertexStartIndex];
				//iVertexIndexBuffer.y = pVertexIndexBuffer[iCurVertexStartIndex+1];
				//iVertexIndexBuffer.z = pVertexIndexBuffer[iCurVertexStartIndex+2];
				int		iCurVertexStartIndex = GetTriangleIndexValue(i)*3;

				iVertexIndexBuffer.x = GetVertexIndexValue(iCurVertexStartIndex);
				iVertexIndexBuffer.y = GetVertexIndexValue(iCurVertexStartIndex+1);
				iVertexIndexBuffer.z = GetVertexIndexValue(iCurVertexStartIndex+2);
				bool bHitTriangles = LineCheck(Start
											, Dir
											, pVertexBuffer[iVertexIndexBuffer.x].Position
											, pVertexBuffer[iVertexIndexBuffer.y].Position
											, pVertexBuffer[iVertexIndexBuffer.z].Position
											, CurNearestHitInfo);

				if(bHitTriangles && CurNearestHitInfo.z<=NearestHitInfo.z)
				{
					NearestHitInfo = CurNearestHitInfo;
					iHitTriangleIndex = i;
				}
			}
		}
		else
		{
			bHit = true;
			//root node need a check
			if(CurProcessNode.NodeIndex!=0)
			{
				CurHitTime = CurProcessNode.HitTime;
			}
			else
			{
				bHit = pCurNode.Bounds.LineCheck(Start, InvDir, CurHitTime);
			}
			if(bHit)
			{
				if(iHitTriangleIndex==-1 || CurHitTime<=NearestHitInfo.z)
				{
					//check child
					bool bHitLeft = pKDTreeRoot[pCurNode.n.Left].Bounds.LineCheck(Start, InvDir, CurHitTimeLeft);
					bool bHitRight = pKDTreeRoot[pCurNode.n.Right].Bounds.LineCheck(Start, InvDir, CurHitTimeRight);
					if(bHitLeft && bHitRight)
					{
						if(CurHitTimeLeft<CurHitTimeRight)
						{
							RootNode.HitTime = CurHitTimeRight;
							RootNode.NodeIndex = pCurNode.n.Right;
							LocalStack.PushElement(RootNode);

							RootNode.HitTime = CurHitTimeLeft;
							RootNode.NodeIndex = pCurNode.n.Left;
							LocalStack.PushElement(RootNode);
						}
						else
						{	
							RootNode.HitTime = CurHitTimeLeft;
							RootNode.NodeIndex = pCurNode.n.Left;
							LocalStack.PushElement(RootNode);

							RootNode.HitTime = CurHitTimeRight;
							RootNode.NodeIndex = pCurNode.n.Right;
							LocalStack.PushElement(RootNode);
						}
					}
					else if(bHitLeft)
					{
						RootNode.HitTime = CurHitTimeLeft;
						RootNode.NodeIndex = pCurNode.n.Left;
						LocalStack.PushElement(RootNode);
					}
					else if(bHitRight)
					{
						RootNode.HitTime = CurHitTimeRight;
						RootNode.NodeIndex = pCurNode.n.Right;
						LocalStack.PushElement(RootNode);
					}
				}
				else
				{
					break;
				}
			}
		}
	}
	UVT = NearestHitInfo;
	return iHitTriangleIndex;
}

_GLOBAL_ void SceneProcessIndirectionalLighting( RasterBufferElement* pRaseter
												, CVertex*	pTriangleVertex
												//, int* pVertexIndexBuffer
												//, int* pTriangleIndexBuffer
												, CKDopNode* pKDTreeRoot
												//, Tuple3f DirectionalLightDir
												, unsigned char* LightMap
												, int* GlobalConstantData
												, Tuple2i LightMapSize
												, int ThreadsCount)
{
	int iCurPixelIndex = blockIdx.x * ThreadsCount + threadIdx.x;
	//current process texelpos
	int Y = iCurPixelIndex/LightMapSize.x;
	int X = iCurPixelIndex - Y*LightMapSize.x;

	RasterBufferElement CurTexel= pRaseter[Y * LightMapSize.x + X];
	unsigned int iCurTexelTriangleIndex = CurTexel.ComponentAndTriangleIndex;
	//LightMap[Y * LightMapSize.x + X] = 0;
	if(iCurTexelTriangleIndex!=InvalidRasterMark)
	{
		//Tuple3i iCurTexelTriangleIdx(pVertexIndexBuffer[iCurTexelTriangleIndex]
		//							,pVertexIndexBuffer[iCurTexelTriangleIndex + 1]
		//							,pVertexIndexBuffer[iCurTexelTriangleIndex + 2]);

		Tuple3i iCurTexelTriangleIdx(GetVertexIndexValue(iCurTexelTriangleIndex)
									,GetVertexIndexValue(iCurTexelTriangleIndex+1)
									,GetVertexIndexValue(iCurTexelTriangleIndex+2));

		Tuple3f V1 = pTriangleVertex[iCurTexelTriangleIdx.x].Position;
		Tuple3f V2 = pTriangleVertex[iCurTexelTriangleIdx.y].Position;
		Tuple3f V3 = pTriangleVertex[iCurTexelTriangleIdx.z].Position;

		Tuple3f CurTexelWorldNormal = pTriangleVertex[iCurTexelTriangleIdx.x].Normal*CurTexel.ElementUV.x
									+ pTriangleVertex[iCurTexelTriangleIdx.y].Normal*CurTexel.ElementUV.y
									+ pTriangleVertex[iCurTexelTriangleIdx.z].Normal*(1.0f-CurTexel.ElementUV.x-CurTexel.ElementUV.y);

		CurTexelWorldNormal = CurTexelWorldNormal.normalize();

		Tuple3f CurTexelWorldPos = V1*CurTexel.ElementUV.x + V2*CurTexel.ElementUV.y + V3 *(1.0f-CurTexel.ElementUV.x-CurTexel.ElementUV.y);
		//
		CurTexelWorldPos.x += CurTexelWorldNormal.x*0.01f;
		CurTexelWorldPos.y += CurTexelWorldNormal.y*0.01f;
		CurTexelWorldPos.z += CurTexelWorldNormal.z*0.01f;

		ProcessIndirectionalLighting(CurTexelWorldPos
									,CurTexelWorldNormal
									,pTriangleVertex
									//,pVertexIndexBuffer
									//,pTriangleIndexBuffer
									,pKDTreeRoot
									,GlobalConstantData
									,&(LightMap[Y * LightMapSize.x + X]));
	}
}

_GLOBAL_ void SceneProcessLighting(RasterBufferElement* pRaseter
									, CVertex*	pTriangleVertex
									//, int* pVertexIndexBuffer
									//, int* pTriangleIndexBuffer
									, CKDopNode* pKDTreeRoot
									, unsigned char* LightMap
									, int* GlobalConstantData
									, Tuple2i LightMapSize
									, int ThreadsCount)
{
	int iCurPixelIndex = blockIdx.x * ThreadsCount + threadIdx.x;
	//current process texelpos
	int Y = iCurPixelIndex/LightMapSize.x;
	int X = iCurPixelIndex - Y*LightMapSize.x;

	RasterBufferElement CurTexel= pRaseter[Y * LightMapSize.x + X];
	unsigned int iCurTexelTriangleIndex = CurTexel.ComponentAndTriangleIndex;
	//LightMap[Y * LightMapSize.x + X] = 0;
	if(iCurTexelTriangleIndex!=InvalidRasterMark)
	{
		/*Tuple3i iCurTexelTriangleIdx(pVertexIndexBuffer[iCurTexelTriangleIndex]
									,pVertexIndexBuffer[iCurTexelTriangleIndex + 1]
									,pVertexIndexBuffer[iCurTexelTriangleIndex + 2]);*/

		Tuple3i iCurTexelTriangleIdx(GetVertexIndexValue(iCurTexelTriangleIndex)
									,GetVertexIndexValue(iCurTexelTriangleIndex+1)
									,GetVertexIndexValue(iCurTexelTriangleIndex+2));

		Tuple3f V1 = pTriangleVertex[iCurTexelTriangleIdx.x].Position;
		Tuple3f V2 = pTriangleVertex[iCurTexelTriangleIdx.y].Position;
		Tuple3f V3 = pTriangleVertex[iCurTexelTriangleIdx.z].Position;

		Tuple3f CurTexelWorldNormal = pTriangleVertex[iCurTexelTriangleIdx.x].Normal*CurTexel.ElementUV.x
									+ pTriangleVertex[iCurTexelTriangleIdx.y].Normal*CurTexel.ElementUV.y
									+ pTriangleVertex[iCurTexelTriangleIdx.z].Normal*(1.0f-CurTexel.ElementUV.x-CurTexel.ElementUV.y);
		
		CurTexelWorldNormal = CurTexelWorldNormal.normalize();

		Tuple3f CurTexelWorldPos = V1*CurTexel.ElementUV.x + V2*CurTexel.ElementUV.y + V3 *(1.0f-CurTexel.ElementUV.x-CurTexel.ElementUV.y);
		//
		CurTexelWorldPos.x += CurTexelWorldNormal.x*0.01f;
		CurTexelWorldPos.y += CurTexelWorldNormal.y*0.01f;
		CurTexelWorldPos.z += CurTexelWorldNormal.z*0.01f;

		float fShadowCoeff = ProcessLightShadow(CurTexelWorldPos
												, CurTexelWorldNormal
												, pTriangleVertex
												//, pVertexIndexBuffer
												//, pTriangleIndexBuffer
												, pKDTreeRoot
												, GlobalConstantData);
		LightMap[Y * LightMapSize.x + X] = fShadowCoeff*255;

		/*ProcessIndirectionalLighting(CurTexelWorldPos
									, CurTexelWorldNormal
									, pTriangleVertex
									, pVertexIndexBuffer
									, pTriangleIndexBuffer
									, pKDTreeRoot
									, GlobalConstantData
									, &(LightMap[Y * LightMapSize.x + X]));*/

		//LightMap[Y * LightMapSize.x + X] = (fShadowCoeff+Irrandiance)*255;

		/*Tuple3f UVT;
		if(SceneLineCheck(CurTexelWorldPos
						, CurTexelWorldNormal
						, UVT
						, pTriangleVertex
						, pVertexIndexBuffer
						, pTriangleIndexBuffer
						, pKDTreeRoot)>=0)
		{
			LightMap[Y * LightMapSize.x + X] = 0;
		}*/

	}
}


 void CBox::AddTriangles(int iStart
						, int iNum
						, const CVertex* pVertexBuffer
						, const int* pVertexIndex
							   ,const int* pTriangleIndex)
{
	for(int i=iStart; i<iStart+iNum; ++i)
	{
		int iTriangleIndex = pTriangleIndex[i]*3;
		AddPoint(pVertexBuffer[pVertexIndex[iTriangleIndex]].Position);
		AddPoint(pVertexBuffer[pVertexIndex[iTriangleIndex+1]].Position);
		AddPoint(pVertexBuffer[pVertexIndex[iTriangleIndex+2]].Position);
	}
	// do some padd space
}

 void CKDopNode::BuildNode(int				Start
						, int	NumTris
						, const CVertex* pVertexBuffer
						, const int*	 pVertexIndex
						, int*	 pTriangleIndex
						, const Tuple3f* pTriangleCenter
						, std::vector<CKDopNode*>&	KDopNodeList)
{
	Bounds.AddTriangles(Start, NumTris, pVertexBuffer, pVertexIndex, pTriangleIndex);
	if(NumTris>MaxTrianglesPerLeaf)
	{
		IsLeaf = 0;
		int BestPlane = -1;
		float BestMean = 0.0f;
		float BestVariance = 0.0f;
		for(int nPlane = 0; nPlane < 3; nPlane++)
		{
			float Mean = 0.0f;
			float Variance = 0.0f;
			for(int nTriangle = Start; nTriangle<Start+NumTris; nTriangle++)
			{
				int iTriangleIndex = pTriangleIndex[nTriangle];
				Mean += pTriangleCenter[iTriangleIndex].getDotProduct(SplitPlane[nPlane]);
			}
			Mean /= float(NumTris);
			for(int nTriangle = Start; nTriangle<Start+NumTris; nTriangle++)
			{
				int iTriangleIndex = pTriangleIndex[nTriangle];
				float Dot = pTriangleCenter[iTriangleIndex].getDotProduct(SplitPlane[nPlane]);
				Variance += (Dot-Mean)*(Dot-Mean);
			}
			Variance /= float(NumTris);
			if(Variance>=BestVariance)
			{
				BestPlane = nPlane;
				BestVariance = Variance;
				BestMean = Mean;
			}
		}
		int Left = Start - 1;
		int Right = Start + NumTris;

		while(Left<Right)
		{
			float Dot;
			int iTriangleIndex;
			do
			{
				iTriangleIndex = pTriangleIndex[++Left];
				Dot = pTriangleCenter[iTriangleIndex].getDotProduct(SplitPlane[BestPlane]);
			}while(Dot<BestMean && Left<Right);

			do
			{
				iTriangleIndex = pTriangleIndex[--Right];
				Dot = pTriangleCenter[iTriangleIndex].getDotProduct(SplitPlane[BestPlane]);
			}while(Dot>=BestMean && Right>0 && Left<Right);

			if(Left<Right)
			{
				int iTemp = pTriangleIndex[Left];
				pTriangleIndex[Left] = pTriangleIndex[Right];
				pTriangleIndex[Right] = iTemp;
			}
		}

		if(Left == Start + NumTris || Right == Start)
		{
			Left = Start + (NumTris/2);
		}

		CKDopNode* TempPadd = new CKDopNode;
		KDopNodeList.push_back(TempPadd);
		n.Left = KDopNodeList.size()-1;

		TempPadd = new CKDopNode;
		KDopNodeList.push_back(TempPadd);
		n.Right = KDopNodeList.size()-1;

		KDopNodeList[n.Left]->BuildNode(  Start
										, Left-Start
										, pVertexBuffer
										, pVertexIndex
										, pTriangleIndex
										, pTriangleCenter
										, KDopNodeList);

		KDopNodeList[n.Right]->BuildNode( Left
										, Start+NumTris-Left
										, pVertexBuffer
										, pVertexIndex
										, pTriangleIndex
										, pTriangleCenter
										, KDopNodeList);
	}
	else
	{
		IsLeaf = 1;
		t.StartTriangleIndex = Start;
		t.TriangleCount = NumTris;
	}
}

CKDop::CKDop(void)
{
	pVertexBuffer = NULL;
	pTriangleIndex = NULL;
	pTriangleCenter = NULL;
	//pKDopNodeTriangleIndex = NULL;
	pd_KDopNodeTriangleIndex = NULL;
	pVertexIndexBuffer = NULL;
	ph_KDopNode = NULL;
	pd_KDopNode = NULL;

	d_TriangleUVInfo = NULL;
	d_TriangleIndex = NULL;
	d_RasterRes = NULL;
	d_TriangleVertex = NULL;


	TextureSize = Tuple2i(512, 512);
	VertexCount = 0;
	VertexIndexCount = 0;
	KDopNodeCount = 0;	
}

CKDop::~CKDop(void)
{
	ReleaseRes();
}

void CKDop::CalcTrianglesCenter()
{
}

void CKDop::BeginBuild()
{
}


void CKDop::BuildKDop()
{
}

void CKDop::EndBuild()
{
}

void CKDop::ReleaseRes()
{
	SAFE_DELETE_ARRAY(pVertexBuffer);
	SAFE_DELETE_ARRAY(pTriangleIndex);
	SAFE_DELETE_ARRAY(pVertexIndexBuffer);
	SAFE_DELETE_ARRAY(pTriangleCenter);
	//SAFE_DELETE_ARRAY(pKDopNodeTriangleIndex);
	SAFE_DELETE_ARRAY(ph_KDopNode);
	SAFE_FREE_CUDA(ph_KDopNode);
	SAFE_FREE_CUDA(pd_KDopNodeTriangleIndex);

	for(int i=0; i<KDopNodes.size(); i++)
	{
		delete KDopNodes[i];
	}
	KDopNodes.clear();
}

void CKDop::Raster()
{
}

void CKDop::PrepareDeviceData()
{
}

void CKDop::ProcessLighting()
{
}

void CSceneManager::PrepareConstantData()
{
	ConstInfoIndex Idx;
	Idx.iRenderConfigOffset = sizeof(ConstInfoIndex)/4;

	int iCurIndex = 0;
	iCurIndex += sizeof(ConstInfoIndex)/sizeof(int);
	memcpy(&hConstantBuffer[iCurIndex], &RenderConfig, sizeof(CRenderConfig));
	iCurIndex += sizeof(CRenderConfig)/sizeof(int);

	Idx.iMeshInfoOffset = iCurIndex;
	//MeshCount
		//MeshVertexIndexStartIndex
		//MeshVertexIndexCount
		//......
	hConstantBuffer[iCurIndex++] = SceneStaticMeshes.size();
	for(int i=0; i<SceneStaticMeshes.size(); ++i)
	{
		hConstantBuffer[iCurIndex++] = SceneStaticMeshes[i]->pInstance->iStartVertexIndex;
		hConstantBuffer[iCurIndex++] = SceneStaticMeshes[i]->pInstance->VertexIndexCount;
		hConstantBuffer[iCurIndex++] = SceneStaticMeshes[i]->pInstance->TextureSize.x;
		hConstantBuffer[iCurIndex++] = SceneStaticMeshes[i]->pInstance->TextureSize.y;
		//memcpy(&hConstantBuffer[iCurIndex], SceneStaticMeshes[i]->MeshEmitColor, sizeof(SceneStaticMeshes[i]->MeshEmitColor));
		//iCurIndex+=sizeof(SceneStaticMeshes[i]->MeshEmitColor)/sizeof(int);
		hConstantBuffer[iCurIndex++] = (int)(d_SceneRasterBufferElement[i]);
	}
	Idx.iLightInfoOffset = iCurIndex;
	//DirectionalLightCount
	hConstantBuffer[iCurIndex++] = DirectionalLights.size();
	//PointLightCount
	hConstantBuffer[iCurIndex++] = PointLights.size();
	//SpotLightCount
	hConstantBuffer[iCurIndex++] = SpotLights.size();
	for(int i=0; i<DirectionalLights.size(); ++i)
	{
		memcpy(&hConstantBuffer[iCurIndex], &DirectionalLights[i], sizeof(CDirectionalLight));
		iCurIndex += sizeof(CDirectionalLight)/sizeof(int);
	}

	for(int i=0; i<PointLights.size(); ++i)
	{
		memcpy(&hConstantBuffer[iCurIndex], &PointLights[i], sizeof(CPointLight));
		iCurIndex += sizeof(CPointLight)/sizeof(int);
	}

	for(int i=0; i<SpotLights.size(); ++i)
	{
		memcpy(&hConstantBuffer[iCurIndex], &SpotLights[i], sizeof(CSpotLight));
		iCurIndex += sizeof(CSpotLight)/sizeof(int);
	}
	Idx.iSampleDirOffset = iCurIndex;

	float sqrt_n_samples = sqrt(float(RenderConfig.SampleDirCount));
	double oneoverN = 1.0/sqrt_n_samples;
	for(int a=0; a<sqrt_n_samples; a++) 
	{
		for(int b=0; b<sqrt_n_samples; b++) 
		{
			// generate unbiased distribution of spherical coords
			double x = (a + rand() * (1.0/RAND_MAX)) * oneoverN; // do not reuse results
			double y = (b + rand() * (1.0/RAND_MAX)) * oneoverN; // each sample must be My_random
			double theta = 2.0 * acos(sqrt(1.0 - x));
			double phi = 2.0 * M_PI * y;
			Tuple3f vec(sin(theta)*cos(phi), sin(theta)*sin(phi), cos(theta));
			vec = vec.normalize();
			memcpy(&hConstantBuffer[iCurIndex], &vec, sizeof(Tuple3f));
			iCurIndex+=sizeof(Tuple3f)/sizeof(int);
		}
	}
	//copy the main index
	memcpy(&hConstantBuffer[0], &Idx, sizeof(ConstInfoIndex));
	//should use constant memory instead!
	//for there are some aligned accesses problems. we use globle memory instead
	cudaMalloc((void**)&pd_ConstantData, MaxConstantBufferSize);
	cudaMemcpy(pd_ConstantData, hConstantBuffer, iCurIndex*sizeof(int), cudaMemcpyHostToDevice);
	//cudaMemcpyToSymbol(dConstantBuffer, hConstantBuffer, iCurIndex*sizeof(int));
}

void CSceneManager::PrepareTextureBuffer()
{
	pVertexIndexBufferTextures[0] = &(VertexIndexBufferTextures_0);
	pVertexIndexBufferTextures[1] = &(VertexIndexBufferTextures_1);
	pVertexIndexBufferTextures[2] = &(VertexIndexBufferTextures_2);
	pVertexIndexBufferTextures[3] = &(VertexIndexBufferTextures_3);
	pVertexIndexBufferTextures[4] = &(VertexIndexBufferTextures_4);
	pVertexIndexBufferTextures[5] = &(VertexIndexBufferTextures_5);
	pVertexIndexBufferTextures[6] = &(VertexIndexBufferTextures_6);
	pVertexIndexBufferTextures[7] = &(VertexIndexBufferTextures_7);
	pVertexIndexBufferTextures[8] = &(VertexIndexBufferTextures_8);
	pVertexIndexBufferTextures[9] = &(VertexIndexBufferTextures_9);


	pTriangleIndexBufferTextures[0] = &(TriangleIndexBufferTextures_0);
	pTriangleIndexBufferTextures[1] = &(TriangleIndexBufferTextures_1);
	pTriangleIndexBufferTextures[2] = &(TriangleIndexBufferTextures_2);
	pTriangleIndexBufferTextures[3] = &(TriangleIndexBufferTextures_3);
	pTriangleIndexBufferTextures[4] = &(TriangleIndexBufferTextures_4);
	pTriangleIndexBufferTextures[5] = &(TriangleIndexBufferTextures_5);
	pTriangleIndexBufferTextures[6] = &(TriangleIndexBufferTextures_6);
	pTriangleIndexBufferTextures[7] = &(TriangleIndexBufferTextures_7);
	pTriangleIndexBufferTextures[8] = &(TriangleIndexBufferTextures_8);
	pTriangleIndexBufferTextures[9] = &(TriangleIndexBufferTextures_9);

	int iNeedVertexIndexTextureCount = VertexIndexCount / Max1DTextureSize;
	int iCurIndex = 0;
	for(int i=0; i<iNeedVertexIndexTextureCount; i++)
	{
		cudaMalloc((void**)&pd_VertexIndexTexture[i] , Max1DTextureSize*sizeof(int));
		cudaMemcpy(pd_VertexIndexTexture[i], &ph_VertexIndexBuffer[Max1DTextureSize*i], Max1DTextureSize*sizeof(int), cudaMemcpyHostToDevice);
		cudaBindTexture(0, *pVertexIndexBufferTextures[i], pd_VertexIndexTexture[i]);
		iCurIndex++;
	}
	iCurIndex--;
	//process last 
	int iLastCount = VertexIndexCount - iNeedVertexIndexTextureCount * Max1DTextureSize;
	if(iLastCount>0)
	{
		iCurIndex++;
		cudaMalloc((void**)&pd_VertexIndexTexture[iCurIndex] , Max1DTextureSize*sizeof(int));
		cudaMemcpy(pd_VertexIndexTexture[iCurIndex], &ph_VertexIndexBuffer[Max1DTextureSize*iCurIndex], iLastCount*sizeof(int), cudaMemcpyHostToDevice);
		cudaBindTexture(0, *pVertexIndexBufferTextures[iCurIndex], pd_VertexIndexTexture[iCurIndex], iLastCount*sizeof(int));
	}
	UsedVertexIndexTextureCount = iCurIndex;

	int iNeedTriangleTextureCount = VertexIndexCount/3 / Max1DTextureSize;
	iCurIndex = 0;
	for(int i=0; i<iNeedTriangleTextureCount; i++)
	{
		cudaMalloc((void**)&pd_TriangleIndexTexture[i] , Max1DTextureSize*sizeof(int));
		cudaMemcpy(pd_TriangleIndexTexture[i], &ph_KDTreeTriangleIndexBuffer[Max1DTextureSize*i], Max1DTextureSize*sizeof(int), cudaMemcpyHostToDevice);
		cudaBindTexture(0, *pTriangleIndexBufferTextures[i], pd_TriangleIndexTexture[i]);	
		iCurIndex++;
	}
	iCurIndex--;
	//process last 
	iLastCount = VertexIndexCount/3 - iNeedTriangleTextureCount * Max1DTextureSize;
	if(iLastCount>0)
	{
		iCurIndex++;
		cudaMalloc((void**)&pd_TriangleIndexTexture[iCurIndex] , Max1DTextureSize*sizeof(int));
		cudaMemcpy(pd_TriangleIndexTexture[iCurIndex], &ph_KDTreeTriangleIndexBuffer[Max1DTextureSize*iCurIndex], iLastCount*sizeof(int), cudaMemcpyHostToDevice);
		cudaBindTexture(0, *pTriangleIndexBufferTextures[iCurIndex], pd_TriangleIndexTexture[iCurIndex]);
	}
	UsedTriangleIndexTextureCount = iCurIndex;

	//VertexIndexBufferTextures_0 = VertexIndexBufferTextures[0];
	//VertexIndexBufferTextures_1 = VertexIndexBufferTextures[1];
	//VertexIndexBufferTextures_2 = VertexIndexBufferTextures[2];
	//VertexIndexBufferTextures_3 = VertexIndexBufferTextures[3];
	//VertexIndexBufferTextures_4 = VertexIndexBufferTextures[4];
	//VertexIndexBufferTextures_5 = VertexIndexBufferTextures[5];
	//VertexIndexBufferTextures_6 = VertexIndexBufferTextures[6];
	//VertexIndexBufferTextures_7 = VertexIndexBufferTextures[7];
	//VertexIndexBufferTextures_8 = VertexIndexBufferTextures[8];
	//VertexIndexBufferTextures_9 = VertexIndexBufferTextures[9];

	//TriangleIndexBufferTextures_0 = TriangleIndexBufferTextures[0];
	//TriangleIndexBufferTextures_1 = TriangleIndexBufferTextures[1];
	//TriangleIndexBufferTextures_2 = TriangleIndexBufferTextures[2];
	//TriangleIndexBufferTextures_3 = TriangleIndexBufferTextures[3];
	//TriangleIndexBufferTextures_4 = TriangleIndexBufferTextures[4];
	//TriangleIndexBufferTextures_5 = TriangleIndexBufferTextures[5];
	//TriangleIndexBufferTextures_6 = TriangleIndexBufferTextures[6];
	//TriangleIndexBufferTextures_7 = TriangleIndexBufferTextures[7];
	//TriangleIndexBufferTextures_8 = TriangleIndexBufferTextures[8];
	//TriangleIndexBufferTextures_9 = TriangleIndexBufferTextures[9];
}

void CSceneManager::ReleaseTextures()
{
	for(int i=0; i<=UsedVertexIndexTextureCount; ++i)
	{
		cudaUnbindTexture(*pVertexIndexBufferTextures[i]);
	}
	for(int i=0; i<=UsedTriangleIndexTextureCount; ++i)
	{
		cudaUnbindTexture(*pTriangleIndexBufferTextures[i]);
	}
}